/*
IDs of data vertices start from 1
IDs of query vertices start from 0
uint *vLabel {
    labels of data vertices, the index is vertex ID and the value is the label of this vertex. vLabel[0] is the total number of vertices
}

map<uint,uint> queryVLabel {
    labels of query vertices. key is vertex ID, value is the label of this vertex. put queryLabels into Phase
}

totWriteRowNum {
    the number of rows writed, which is new emb. use cudaMemcpyAsync to initialize totWriteRowNum_dev into 0.
}

class recordPosInfo {
    uint size;
    uint *addr;
} {
    size is the number of uints that addr points to.
    addr is positions of each source vid belong to the specified key. each position contains two uints, the first 
    is the distance between the specified evid label position (including index vertices) and the starting address of 
    neighbors of this svid in this edge label partition. the second is the len of neighbors (not including index vertices).
    if the firsth uint is 1 and the second is 0, there is no neighbors belong to the evid 
    label for this svid under the specified edge label. if the first is 0 and the second is 0, we need to test if there 
    are neighbors belong to the evid label for this svid under the specified edge label. if the first >=0, second >0, we 
    have the address. All first and second uints are initialized to 0
}

map<uint,recordPosInfo> recordPosData {
    the key is combination of edge label and ext vid label. edge label is high 32 bits, ext vid label is low 32 bits.
    the value is the recordPosInfo. size is the number of src vertices in the corresponding edge label partition.
    addr is the address of allocateed space in host memory for the key. Noted that the space allocated for the keys that
    have the same edge label are contiguous. the special key is edge label is high 32 bits and 0xffffffff is low 32 bits.
    this means all keys that have the same edge label. the size in the corresponding recordPosInfo is the tot src vertices
    number belong to the same edge label. the addr in the corresponding recordPosInfo is the starting address of this contiguous
    allocated space.
}

vector<uint> auxArray {
    for reduction:
    this data is used for reduction to store svid and evid of reduction edges and other information. The format is
    auxArray[0] = totUintNum, the number uints that used to stroe positions of svid and evid of reductions. each position
        takes 5 bits, therefore, each uint can hold 3 pairs. not include auxArray[0], we at most have 16 pairs (6 uints),
        therefore, the storage is fiexed to 6+1.
    auxArray[1] = { svid1pos,evid1pos,svid2pos,evid2pos,svid3pos,evid3pos},...,. auxArray[1+totUintNum] = evid1label,
    auxArray[1+totUintNum+1] = evid2label, auxArray[1+totUintNum+2] = evid3label,..., auxArray[1+totUintNum+16] = evid16label (if has).
    we use fixed 16 uints to store labels. auxArray[1+totUintNum+16+1] = {record pos of this evidlabel},..., we use fixed 16 uints
    to store record pos of the correspoing evidlabel. auxArray[1+totUintNum+16+16] = recordFlag, low 16 bits are useRecord, high 16 bits
    are isRecord;
    | totuint  | vid pos  |evid labels| record pos|   flag   |
    |--1 uint--|--6 uint--|--16 uint--|--16 uint--|--1 uint--|

    for extension:
    this data is used to store positions of vertices that has the same vertex label with extvid1 and extvid2, and other information.
    auxArray[0]=evid1label,auxArray[1]=evid2label,auxArray[2]=pos of restrict vertex of evid1, auxArray[3]=pos of restrict vertex of evid2.
    auxArray[4]=pos of svid1,auxArray[5]=pos of svid2.
    auxArray[6] is the number of vertices that has the same vertex label with extvid1. 
    auxArray[7] is the total number of vertices that has the same vertex albel with extvid1 and extvid2
    auxArray[8],..., are the positions of equal vertices for extvid1 and extvid2
}

map<uint,vector<uint>> equalPos {
    key is the vertex label, value is positions of query vertices that have the same label as key.
}

map<uint64_t,uint> repeatFlag {
    key is the combination of edge label and evid label, edge label is high 32 bits, evid label is low 32 bits. value is a combination
    of two 16 bit numbers, the high 16 bits represent how many times this repeated and this value is not changed during mining.
    the low 16 bits represents how many times left for this key.
}

map<uint,uint> restricts {
    if (vid1,vid2) has restrict, then key is vid2, value is vid1
}

map<uint,uint*> allEdgeLabelPartitions {
    key is the edge label in this data base, value is the address of the corresponding edge label partition. for the key 0xffffffff, the value
    is vLabel. the format of edge label partition edgeLabelPartition is
    |--uint0--|--uint1--|--uint2--|--uint3--|--uint4--|--array0--|--array1--|--array2--|--array3--|
    array0: interval data, at most contains 256 intervals. each interval has two uints, the first uint is the starting vid of this 
            interval, the second uint is the number of vertices that this and all previous intervals represent. if there are n intervals,
            the first n uints in array0 are starting vids of all intervals, array0[n] is 0, array0[n+1] is the number of vertices in
            the first interval, array0[n+2] is the number of vertices in the first and second intervals. the last element of array0 is the 
            svid that separates array0 and array2
    array1: the distance between the starting address of array3 and the starting address of neigbors of the corresponding interval vid.
            need extra uint to record the end address for calculating the length of neighbors of last interval vertex
    array2: hash index
    array3: neighbors, neighbors of a vertex are grouped by end vertex label, the labels are in increasing order,
    uint0: total uints in this edge label partition. including uint0
    uint1: the number of source vids in array1 plus (the number of uints in array2)/2, because array2 needs to record vid and postion. we only need one
            uint in our recordpos. this is used for recordPos, we do not count the extra uint in array1.
    uint2: the number of intervals
    uint3: the distance between the starting address of array2 and the starting address of array0
    uint4: the distance between the starting address of array3 and the starting address of array0

}

vector<Phase> matchOrder{
    if the phase is extension, and the one srouce vid and two evids are not the same label.
    ext1label < ext2label
}

gpu memoyr arrangement
|--vertex label--|--totWriteRowNum--|--edgeLabelPartition--|--recordPos(if has)--|--auxArray(if has)--|--new emb--|--partial emb--|


basic naming conventions of arguments{
    uint *baseAddr_dev, indicates the starting address of allocated gpu memory with vertex label and totWriteRowNum skipped
    uint *partialEmb, generted embeddings;    
    uint *edgeLabelPartition, edge label partition for the specified edge label.
    uint *edgeLabelPartition_dev, stores elements of array0, array1, array2 and array3 in device memory
    uint *neigborsData_dev, the starting address of array3
    uint totSrcNum, uint1
    uint intervalNum, uint2
    uint *totWriteRowNum_dev, record how many rows are writed in this phase, the number of generated rows in the new embedding
    uint partialRowNum, the number of rows in the partialEmb
}

vector<queryNeigInfo> query{
    each query vertex has an entry in query. 
    query[vid].neigs is the neighbors of vid in query.
    since we have at most 32 query vertices, we use an uint to represent neighbors, each bit
    represents a neighbor.
    query[vid].edgeLabelNeigs is map<uint,vector<uint>>. this is the neigbors grouped by edge
    labels, the key of the map is the edge label. and the vector is neighbors. note that 
    vector[0] is the number of neighbors belong to this group, vector[1] is the bit represention
    of neighbors 
}
*/


#include <iostream>
#include <fstream>
#include <cuda_runtime.h>
#include <algorithm>
#include <hash_map>
#include <ctime>
#include "common.h"
#include "loadDataBase.h"
#include "genMatchOrder.h"
#include "miningPattern.h"
#include "reduction.h"

using namespace __gnu_cxx;
class recordPosInfo{
public:
    uint size;
    uint* addr;
};
size_t gpuAvailSpace;
uint maxRowNum;
std::vector<uint*> partResults;
std::vector<uint> partResultRowNum;
int GPU_SM_NUM;
int debugNum=0;
int debugNumend=0;
std::map<uint,uint> posOfVisitedV;
std::map<uint64_t,recordPosInfo> recordPosData;
std::map<uint,std::vector<uint>> equalPos;
std::vector<Phase> matchOrder;
std::map<uint64_t,uint> repeatFlag;
std::map<uint,std::map<uint,bool>> edgeLabelStat;
std::map<uint,uint> restricts;
std::map<uint,uint*> allEdgeLabelPartitions;
std::map<uint,uint> queryLabel;
uint *vLabel;


void printPartialEmbGPUDebug(uint *partialEmb_dev, uint &partialEmbRowNum, uint embLen, std::string filename){
    //partialEmbRowNum = debugNumend-debugNum;
    uint *partialEmb_host = (uint*)malloc(sizeof(uint)*partialEmbRowNum*embLen);
    std::ifstream inputfile(filename,std::ios::in);
    for(uint i=0;i<partialEmbRowNum;++i){
        for(uint j=0;j<embLen;++j){
            inputfile>>partialEmb_host[i*embLen+j];
        }
    }
    cudaMemcpy(partialEmb_dev,partialEmb_host,sizeof(uint)*partialEmbRowNum*embLen,cudaMemcpyHostToDevice);
    inputfile.close();
    free(partialEmb_host);
}
void printPartialEmbGPU(uint *partialEmb_host, uint partialEmbRowNum, uint embLen, std::ofstream &outputfile){
    for(uint i=0;i<partialEmbRowNum;++i){
        outputfile<<partialEmb_host[i*embLen+0];
        for(uint j=1;j<embLen;++j){
            outputfile<<" "<<partialEmb_host[i*embLen+j];
        }
        outputfile<<std::endl;
    }
}
void printPartialEmbGPU(uint *partialEmb_dev, uint partialEmbRowNum, uint embLen, std::ofstream &outputfile, int tmp){
    uint *partialEmb_host = (uint*)malloc(sizeof(uint)*partialEmbRowNum*embLen);
    cudaMemcpy(partialEmb_host,partialEmb_dev,sizeof(uint)*partialEmbRowNum*embLen,cudaMemcpyDeviceToHost);
    for(uint i=0;i<partialEmbRowNum;++i){
        outputfile<<partialEmb_host[i*embLen+0];
        for(uint j=1;j<embLen;++j){
            outputfile<<" "<<partialEmb_host[i*embLen+j];
        }
        outputfile<<std::endl;
    }
    free(partialEmb_host);
}
void printPartialEmbGPU(uint *partialEmb_dev, uint partialEmbRowNum, uint embLen, std::string filename){
    uint *partialEmb_host = (uint*)malloc(sizeof(uint)*partialEmbRowNum*embLen);
    cudaMemcpy(partialEmb_host,partialEmb_dev,sizeof(uint)*partialEmbRowNum*embLen,cudaMemcpyDeviceToHost);
    std::ofstream outputfile;
    outputfile.open(filename, std::ios::out);
    if (!outputfile.is_open()) return;
    for(uint i=0;i<partialEmbRowNum;++i){
        outputfile<<partialEmb_host[i*embLen+0];
        for(uint j=1;j<embLen;++j){
            outputfile<<" "<<partialEmb_host[i*embLen+j];
        }
        outputfile<<std::endl;
    }
    outputfile.close();
    free(partialEmb_host);
}
void constructRecordPos(){
    for(auto ite=edgeLabelStat.begin();ite!=edgeLabelStat.end();++ite){
        uint edgelabel = ite->first;
        uint *tmplabelpart = allEdgeLabelPartitions[edgelabel];
        uint vertexNum = tmplabelpart[1];
        uint typeNum = ite->second.size();
        uint *base = new uint[vertexNum*typeNum*2]();
        uint64_t key = edgelabel;
        key = (key<<32)|0xffffffff;
        recordPosInfo tmp;
        tmp.size = vertexNum*typeNum*2;
        tmp.addr = base;
        recordPosData.insert(std::pair<uint64_t,recordPosInfo>(key,tmp));
        for(auto itea=ite->second.begin();itea!=ite->second.end();++itea){
            uint vlabel = itea->first;
            uint64_t keya = edgelabel;
            keya=(keya<<32)|vlabel;
            recordPosInfo tmpa;
            tmpa.size = vertexNum*2;
            tmpa.addr = base;
            recordPosData.insert(std::pair<uint64_t,recordPosInfo>(keya,tmpa));
            base = base + vertexNum*2;
        }
    }
}

uint genInit(uint *vLabel_dev, uint *writeRowNum_dev, uint *edgeLabelPartition_dev, uint *neighborsData_dev, 
    uint *baseAddr_dev, uint stride, uint intervalNum, uint totSrcNum,bool isContinue, uint blocksize){

    Phase &firstPhase = matchOrder[0];
    uint edgeLabel = firstPhase.edgeLabel;
    std::vector<uint> &pairs = firstPhase.pairs;
    if(pairs.size()==2){
        uint svid = pairs[0], evid = pairs[1];
        posOfVisitedV.insert(std::pair<uint,uint>(svid,0));
        posOfVisitedV.insert(std::pair<uint,uint>(evid,1));
        uint svidlabel = queryLabel[svid], evidlabel = queryLabel[evid];
        uint64_t key = edgeLabel;
        key = (key<<32)|evidlabel;
        bool isExtRestrict = restricts[evid]==svid;
        if(evidlabel==svidlabel){
            std::vector<uint> tmpv;
            tmpv.push_back(0);
            tmpv.push_back(1);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svidlabel,tmpv));
        }else{
            std::vector<uint> tmpv1;
            std::vector<uint> tmpv2;
            tmpv1.push_back(0);
            tmpv2.push_back(1);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svidlabel,tmpv1));
            equalPos.insert(std::pair<uint,std::vector<uint>>(evidlabel,tmpv2));
        }

        //the generated rows can not occupy half of the available sapec
        maxRowNum = ((gpuAvailSpace-stride)/2-38*1500)/2;
        if(intervalNum<=1024){
            //the last parameter is levelIndexNum
            initEmb_2V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,baseAddr_dev+stride,intervalNum,totSrcNum,svidlabel,evidlabel,isExtRestrict,isContinue,maxRowNum,0);
        }else{
            std::cout<<"before initEmb_2V_NoHash intervalNum>1024"<<std::endl;
            initEmb_2V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,baseAddr_dev+stride,intervalNum,totSrcNum,svidlabel,evidlabel,isExtRestrict,isContinue,maxRowNum,blocksize);
            cudaDeviceSynchronize();
            cudaError_t err = cudaGetLastError();
            std::cout<<"################"<<cudaGetErrorString(err)<<std::endl;
        }
    }else{
        std::cout<<"in the init_3V"<<std::endl;
        uint svid1 = pairs[0],evid1 = pairs[1],svid2 = svid1,evid2 = pairs[3];
        uint svid1label = queryLabel[svid1], evid1label = queryLabel[evid1], svid2label = svid1label, evid2label = queryLabel[evid2];
        uint64_t key1 = edgeLabel, key2=edgeLabel;
        key1 = (key1<<32)|evid1label; key2 = (key2<<32)|evid2label;
        posOfVisitedV.insert(std::pair<uint,uint>(svid1,0)); posOfVisitedV.insert(std::pair<uint,uint>(evid1,1)); posOfVisitedV.insert(std::pair<uint,uint>(evid2,2));
        bool isExt1Restrict = restricts[evid1]!=0xffffffff;
        bool isExt1and2Restrict = restricts[evid2]!=0xffffffff?restricts[evid2]==evid1:false;
        bool isExt2Restrict = isExt1and2Restrict?false:restricts[evid2]!=0xffffffff;
        bool isExt1and2SameLabel = evid1label==evid2label;
        if(svid1label==evid1label && svid1label==evid2label){
            std::vector<uint> tmpv;
            tmpv.push_back(0); tmpv.push_back(1); tmpv.push_back(2);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svid1label,tmpv));
        }else if(svid1label==evid1label){
            std::vector<uint> tmpv1; std::vector<uint> tmpv2;
            tmpv1.push_back(0); tmpv1.push_back(1); tmpv2.push_back(2);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svid1label,tmpv1)); equalPos.insert(std::pair<uint,std::vector<uint>>(evid2label,tmpv2));
        }else if(svid1label==evid2label){
            std::vector<uint> tmpv1; std::vector<uint> tmpv2;
            tmpv1.push_back(0); tmpv1.push_back(2); tmpv2.push_back(1);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svid1label,tmpv1)); equalPos.insert(std::pair<uint,std::vector<uint>>(evid1label,tmpv2));
        }else if(evid1label==evid2label){
            std::vector<uint> tmpv1; std::vector<uint> tmpv2;
            tmpv1.push_back(0); tmpv2.push_back(1); tmpv2.push_back(2);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svid1label,tmpv1)); equalPos.insert(std::pair<uint,std::vector<uint>>(evid1label,tmpv2));
        }else{
            std::vector<uint> tmpv1; std::vector<uint> tmpv2; std::vector<uint> tmpv3;
            tmpv1.push_back(0); tmpv2.push_back(1); tmpv3.push_back(2);
            equalPos.insert(std::pair<uint,std::vector<uint>>(svid1label,tmpv1)); equalPos.insert(std::pair<uint,std::vector<uint>>(evid1label,tmpv2)); equalPos.insert(std::pair<uint,std::vector<uint>>(evid2label,tmpv3));
        }
        uint *recordPos1_dev = NULL, *recordPos2_dev = NULL;
        maxRowNum = ((gpuAvailSpace-stride)/2-38*1500)/3;
        //maxRowNum = 18000;
        std::cout<<"before init_3V"<<std::endl;
        if(intervalNum<=1024){
            initEmb_3V_1and2_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,baseAddr_dev+stride,intervalNum,totSrcNum,
            svid1label,evid1label,evid2label,isExt1Restrict,isExt2Restrict,isExt1and2Restrict,isExt1and2SameLabel,isContinue,maxRowNum,0);
        }else{
            initEmb_3V_1and2_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,baseAddr_dev+stride,intervalNum,totSrcNum,
            svid1label,evid1label,evid2label,isExt1Restrict,isExt2Restrict,isExt1and2Restrict,isExt1and2SameLabel,isContinue,maxRowNum,blocksize);
        }
        cudaDeviceSynchronize();
        cudaError_t err = cudaGetLastError();
        std::cout<<"after initemb_3v "<<cudaGetErrorString(err)<<std::endl;
    }
    return stride;
}

uint* genExtEmb(uint *vLabel_dev, uint *writeRowNum_dev, uint *writeRowNum, uint *edgeLabelPartition_dev, uint *neighborsData_dev,
    uint *baseAddr_dev, uint &stride, uint *partialEmb_dev, uint intervalNum, uint totSrcNum, uint partialRowNum, 
    Phase &curPhase, bool isLastPhase, bool isContinue, uint blocksize){
	cudaError_t err;
    uint edgeLabel = curPhase.edgeLabel;
    std::vector<uint> &pairs = curPhase.pairs;
    if(pairs.size()==2){
        uint svid = pairs[0], evid = pairs[1];
        uint index = posOfVisitedV.size();
        uint preEmbLen = posOfVisitedV.size();
        posOfVisitedV.insert(std::pair<uint,uint>(evid,index));
        uint embLen = posOfVisitedV.size();
        uint svidlabel = queryLabel[svid], evidlabel = queryLabel[evid];
        uint64_t key = edgeLabel;
        key = (key<<32)|evidlabel;
        bool isExtRestrict = restricts[evid]!=0xffffffff;

        std::vector<uint> auxArray(8);
        auxArray[0]=evidlabel;auxArray[1]=evidlabel;
        auxArray[4]=posOfVisitedV[svid];auxArray[5]=posOfVisitedV[svid];
        auxArray[2]=isExtRestrict?posOfVisitedV[restricts[evid]]:0;
        auxArray[3]=auxArray[2];
        if(equalPos.find(evidlabel)!=equalPos.end()){
            std::vector<uint> &tmpv = equalPos[evidlabel];
            auxArray[6]=tmpv.size();
            auxArray[7]=tmpv.size();
            auxArray.insert(auxArray.end(),tmpv.begin(),tmpv.end());
            tmpv.push_back(posOfVisitedV[evid]);
        }else{
            auxArray[5]=0;
            auxArray[6]=0;
            std::vector<uint> tmpv;
            tmpv.push_back(posOfVisitedV[evid]);
            equalPos.insert(std::pair<uint,std::vector<uint>>(evidlabel,tmpv));
        }
        uint *auxArray_dev = baseAddr_dev+stride;
        stride=stride+auxArray.size();
        cudaMemcpyAsync(auxArray_dev,auxArray.data(),sizeof(uint)*auxArray.size(),cudaMemcpyHostToDevice);
        uint lesspos = isExtRestrict?posOfVisitedV[restricts[evid]]:0;
        std::cout << "before extEmb_1V_NoHash"<< std::endl;
        maxRowNum = (gpuAvailSpace-stride-partialRowNum*preEmbLen-1500*38)/embLen;
        //maxRowNum = 30000;
        if(intervalNum<=1024){
            extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,false,maxRowNum,0);
        }else{
            extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,false,maxRowNum,blocksize);
        }
        //cudaDeviceSynchronize();
        cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
        uint tmppartialRowNum = writeRowNum[0];
        std::cout<<"extEmb_1V_NoHash is done: partialRowNum="<<tmppartialRowNum<<" max="<<maxRowNum<<std::endl;
        //printPartialEmbGPU(baseAddr_dev+stride, tmppartialRowNum,embLen,"testext1.txt");
	    /*if(tmppartialRowNum>=maxRowNum){
	    	std::cout<<"extEmb_1V_NoHash exceeds the maximum: genrownum="
	    		 <<tmppartialRowNum<<" maxrownum="<<maxRowNum<<" remove="
	    		 <<writeRowNum[1]<<std::endl;
	    }*/
        std::vector<uint*> tmppartResults;
        std::vector<uint> tmppartResultRowNum;
        while(tmppartialRowNum>=maxRowNum){
            tmppartialRowNum = tmppartialRowNum-writeRowNum[1];
            std::cout<<"real partialRowNum="<<tmppartialRowNum<<std::endl;
            uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
            uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
            cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            tmppartResults.push_back(tmpaddr);
            tmppartResultRowNum.push_back(tmppartialRowNum);
            //cudaDeviceSynchronize();
            //err = cudaGetLastError();
            //std::cout<<"in extension 2 "<<cudaGetErrorString(err)<<" before extEmb_1V"<<std::endl;
            std::cout<<"continue extEmb_1V_NoHash"<<std::endl;
            if(intervalNum<=1024){
                extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                    baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,true,maxRowNum,0);
            }else{
                extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                    baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,true,maxRowNum,blocksize);
            }
           
            cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
            tmppartialRowNum = writeRowNum[0];
            //err = cudaGetLastError();
            //std::cout<<"in extension 2 "<<cudaGetErrorString(err)<<" after extEmb_1V:"<<tmppartialRowNum<<std::endl;
        }
        std::cout<<"real partialRowNum="<<tmppartialRowNum<<" end continue"<<std::endl;
        if(partResults.size()>0){
            std::cout<<"there are other parts to be processed"<<std::endl;
            uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
            uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
            cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            tmppartResults.push_back(tmpaddr);
            tmppartResultRowNum.push_back(tmppartialRowNum);
            for(uint j=0;j<partResults.size();++j){
                tmpaddr = partResults[j];
                partialRowNum = partResultRowNum[j];
                partialEmb_dev = baseAddr_dev+gpuAvailSpace-partialRowNum*preEmbLen-32;
                cudaMemcpyAsync(partialEmb_dev,tmpaddr,sizeof(uint)*partialRowNum*preEmbLen,cudaMemcpyHostToDevice);
                maxRowNum = (gpuAvailSpace-stride-partialRowNum*preEmbLen-1500*38)/embLen;
                if(intervalNum<=1024){
                    extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                        baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,false,maxRowNum,0);
                }else{
                    extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                        baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,false,maxRowNum,blocksize);
                }

                cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
                free(tmpaddr);
                tmppartialRowNum = writeRowNum[0];
                while(tmppartialRowNum>=maxRowNum){
                    tmppartialRowNum = tmppartialRowNum-writeRowNum[1];
                    uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
                    uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
                    cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
                    writeRowNum[0] = 0;
                    writeRowNum[1] = 0;
                    cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
                    tmppartResults.push_back(tmpaddr);
                    tmppartResultRowNum.push_back(tmppartialRowNum);
                    if(intervalNum<=1024){
                        extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                            baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,true,maxRowNum,0);
                    }else{
                        extEmb_1V_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,auxArray_dev,
                            baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,isExtRestrict,isLastPhase,true,maxRowNum,blocksize);
                    }

                    cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
                    tmppartialRowNum = writeRowNum[0];
                }
            }
        }
        partResults = tmppartResults;
        partResultRowNum = tmppartResultRowNum;
    }else{
        uint svid1 = pairs[0],evid1 = pairs[1],svid2 = pairs[2],evid2 = pairs[3];
        uint index = posOfVisitedV.size();
        uint preEmbLen = posOfVisitedV.size();
        posOfVisitedV.insert(std::pair<uint,uint>(evid1,index));
        posOfVisitedV.insert(std::pair<uint,uint>(evid2,index+1));
        uint embLen = posOfVisitedV.size();
        uint svid1label = queryLabel[svid1],evid1label = queryLabel[evid1];
        uint svid2label = queryLabel[svid2],evid2label = queryLabel[evid2];
        uint64_t key1 = edgeLabel, key2 = edgeLabel;
        key1 = (key1<<32)|evid1label; 
        key2 = (key2<<32)|evid2label;

        bool isExt1Restrict = restricts[evid1]!=0xffffffff;
        bool isExt1and2Restrict = restricts[evid2]!=0xffffffff?restricts[evid2]==evid1:false;
        bool isExt2Restrict = isExt1and2Restrict?false:restricts[evid2]!=0xffffffff;
        bool isExt1and2SameLabel = evid1label==evid2label;
        bool isExt1and2SameSrc = svid1==svid2;

        uint *recordPos1_dev = NULL, *recordPos2_dev=NULL;

        std::vector<uint> auxArray(8);
        auxArray[0]=evid1label;auxArray[1]=evid2label;auxArray[4]=posOfVisitedV[svid1];auxArray[5]=posOfVisitedV[svid2];
        auxArray[2] = isExt1Restrict?posOfVisitedV[restricts[evid1]]:0;
        auxArray[3] = isExt2Restrict?posOfVisitedV[restricts[evid2]]:0;
        if(equalPos.find(evid1label)!=equalPos.end()){
            std::vector<uint> &tmpv = equalPos[evid1label];
            auxArray[6]=tmpv.size();
            auxArray.insert(auxArray.end(),tmpv.begin(),tmpv.end());
            uint index = posOfVisitedV[evid1];
            tmpv.push_back(index);
        }else{
            std::vector<uint> tmpv;
            uint index = posOfVisitedV[evid1];
            tmpv.push_back(index);
            equalPos.insert(std::pair<uint,std::vector<uint>>(evid1label,tmpv));
            auxArray[6]=0;
        }
        if(!isExt1and2SameLabel){
            if(equalPos.find(evid2label)!=equalPos.end()){
                std::vector<uint> &tmpv = equalPos[evid2label];
                uint tmpsize = tmpv.size();
                auxArray[7]=tmpsize+auxArray[6];
                auxArray.insert(auxArray.end(),tmpv.begin(),tmpv.end());
                uint index = posOfVisitedV[evid2];
                tmpv.push_back(index);
            }else{
                auxArray[7]=0+auxArray[6];
                std::vector<uint> tmpv;
                uint index = posOfVisitedV[evid2];
                tmpv.push_back(index);
                equalPos.insert(std::pair<uint,std::vector<uint>>(evid2label,tmpv));
            }
        }else{
            auxArray[7] = auxArray[6];
            std::vector<uint> &tmpv = equalPos[evid1label];
            uint index = posOfVisitedV[evid2];
            tmpv.push_back(index);
        }

        uint *auxArray_dev = baseAddr_dev+stride;
        stride=stride+auxArray.size();
        cudaMemcpyAsync(auxArray_dev,auxArray.data(),sizeof(uint)*auxArray.size(),cudaMemcpyHostToDevice);
        maxRowNum = (gpuAvailSpace-stride-partialRowNum*preEmbLen-1500*38)/embLen;
        if(intervalNum<=1024){
            if(isExt1and2SameSrc && isExt1and2SameLabel) {
                std::cout << "before extEmb_2V_1src_1and2_sameLabel_NoHash"<< std::endl;
                extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                    isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,0);
            }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                std::cout << "before extEmb_2V_1src_1and2_notSameLabel_NoHash"<< std::endl;
                extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                    embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,0);
            }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                std::cout << "before extEmb_2V_2src_1and2_sameLabel_NoHash"<< std::endl;
                extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                    isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,0);
            }else{
                std::cout << "before extEmb_2V_2src_1and2_notSameLabel_NoHash"<< std::endl;
                extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                    embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,0);
            }
        }else{
            if(isExt1and2SameSrc && isExt1and2SameLabel) {
                std::cout << "before extEmb_2V_1src_1and2_sameLabel_NoHash"<< std::endl;
                extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                    isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,blocksize);
            }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                std::cout << "before extEmb_2V_1src_1and2_notSameLabel_NoHash"<< std::endl;
                extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                    embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,blocksize);
            }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                std::cout << "before extEmb_2V_2src_1and2_sameLabel_NoHash"<< std::endl;
                extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                    isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,blocksize);
            }else{
                std::cout << "before extEmb_2V_2src_1and2_notSameLabel_NoHash"<< std::endl;
                extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                    auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                    embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,blocksize);
            }
        }
        cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
        uint tmppartialRowNum = writeRowNum[0];
        std::cout<<"extEmb_2V is done: partialRowNum="<<tmppartialRowNum<<" max="<<maxRowNum<<std::endl;
        std::vector<uint*> tmppartResults;
        std::vector<uint> tmppartResultRowNum;
        while(tmppartialRowNum>=maxRowNum){
            //std::cout<<"over size detec in extension 4"<<std::endl;
            tmppartialRowNum = tmppartialRowNum-writeRowNum[1];
            std::cout<<"real partialRowNum="<<tmppartialRowNum<<std::endl;
            uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
            uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
            cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            tmppartResults.push_back(tmpaddr);
            tmppartResultRowNum.push_back(tmppartialRowNum);
            std::cout<<"continue extEmb_2V_..._NoHash"<<std::endl;
            if(intervalNum<=1024){
                if(isExt1and2SameSrc && isExt1and2SameLabel) {
                    extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                        isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,0);
                }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                    extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                        embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,0);
                }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                    extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                        isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,0);
                }else{
                    extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                        embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,0);
                }
            }else{
                if(isExt1and2SameSrc && isExt1and2SameLabel) {
                    extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                        isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,blocksize);
                }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                    extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                        embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,blocksize);
                }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                    extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                        isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,blocksize);
                }else{
                    extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                        auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                        embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,blocksize);
                }
            }
            cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
            tmppartialRowNum = writeRowNum[0];
        }
        std::cout<<"real partialRowNum="<<tmppartialRowNum<<" end continue"<<std::endl;
        if(partResults.size()>0){
            uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
            uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
            cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            tmppartResults.push_back(tmpaddr);
            tmppartResultRowNum.push_back(tmppartialRowNum);
            for(uint j=0;j<partResults.size();++j){
                tmpaddr = partResults[j];
                partialRowNum = partResultRowNum[j];
                partialEmb_dev = baseAddr_dev+gpuAvailSpace-partialRowNum*preEmbLen-32;
                cudaMemcpyAsync(partialEmb_dev,tmpaddr,sizeof(uint)*partialRowNum*preEmbLen,cudaMemcpyHostToDevice);
                maxRowNum = (gpuAvailSpace-stride-partialRowNum*preEmbLen-1500*38)/embLen;
                if(intervalNum<=1024){
                    if(isExt1and2SameSrc && isExt1and2SameLabel) {
                        extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                            isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,0);
                    }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                        extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                            embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,0);
                    }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                        extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                            isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,0);
                    }else{
                        extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                            embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,0);
                    }
                }else{
                    if(isExt1and2SameSrc && isExt1and2SameLabel) {
                        extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                            isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,blocksize);
                    }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                        extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                            embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,blocksize);
                    }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                        extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                            isExt1Restrict,isExt2Restrict,isExt1and2Restrict,false,maxRowNum,blocksize);
                    }else{
                        extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                            auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                            embLen,isExt1Restrict,isExt2Restrict,false,maxRowNum,blocksize);
                    }
                }

                cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
                free(tmpaddr);
                tmppartialRowNum = writeRowNum[0];
                while(tmppartialRowNum>=maxRowNum){
                    tmppartialRowNum = tmppartialRowNum-writeRowNum[1];
                    uint allocnum = ((tmppartialRowNum*embLen+31)>>5)<<5;
                    uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
                    cudaMemcpyAsync(tmpaddr,baseAddr_dev+stride,sizeof(uint)*tmppartialRowNum*embLen,cudaMemcpyDeviceToHost);
                    writeRowNum[0] = 0;
                    writeRowNum[1] = 0;
                    cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
                    tmppartResults.push_back(tmpaddr);
                    tmppartResultRowNum.push_back(tmppartialRowNum);
                    if(intervalNum<=1024){
                        if(isExt1and2SameSrc && isExt1and2SameLabel) {
                            extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                                isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,0);
                        }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                            extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                                embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,0);
                        }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                            extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                                isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,0);
                        }else{
                            extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                                embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,0);
                        }
                    }else{
                        if(isExt1and2SameSrc && isExt1and2SameLabel) {
                            extEmb_2V_1src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                                isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,blocksize);
                        }else if(isExt1and2SameSrc && !isExt1and2SameLabel){
                            extEmb_2V_1src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                                embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,blocksize);
                        }else if(!isExt1and2SameSrc && isExt1and2SameLabel) {
                            extEmb_2V_2src_1and2_sameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,embLen,
                                isExt1Restrict,isExt2Restrict,isExt1and2Restrict,true,maxRowNum,blocksize);
                        }else{
                            extEmb_2V_2src_1and2_notSameLabel_NoHash(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,
                                auxArray_dev,baseAddr_dev+stride,partialEmb_dev,intervalNum,partialRowNum,
                                embLen,isExt1Restrict,isExt2Restrict,true,maxRowNum,blocksize);
                        }
                    }

                    cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
                    tmppartialRowNum = writeRowNum[0];
                }
            }
        }
        partResults = tmppartResults;
        partResultRowNum = tmppartResultRowNum;

    }
    return partialEmb_dev;
}

void sortReducePairs(std::vector<uint>& pairs){
    uint size = pairs.size()/2;
    for(int i=0;i<size;++i){
        uint tmpsize = size-i;
        uint minsvid = pairs[0];
        uint minevid = pairs[1];
        uint minindex=0;
        for(int j=1;j<tmpsize;++j){
            uint cursvid = pairs[j*2];
            uint curevid = pairs[j*2+1];
            if(cursvid<minsvid){
                minsvid = cursvid;
                minevid = curevid;
                minindex = j;
            }else if(cursvid==minsvid){
                if(queryLabel[cursvid]<queryLabel[minsvid]){
                    minsvid = cursvid;
                    minevid = curevid;
                    minindex = j;
                }else if(queryLabel[curevid]==queryLabel[minevid]){
                    if(curevid<minevid){
                        minsvid = cursvid;
                        minevid = curevid;
                        minindex = j;
                    }
                }
            }
        }
        uint lastindex = tmpsize-1;
        uint tmp = pairs[minindex*2];
        pairs[minindex*2] = pairs[lastindex*2];
        pairs[lastindex*2] = tmp;
        tmp = pairs[minindex*2+1];
        pairs[minindex*2+1] = pairs[lastindex*2+1];
        pairs[lastindex*2+1] = tmp;
    }
}

//the format for indexForIndex[0:2*256] is vs1,vs2,vs3,len1,len2,len3. len1 is the numberr of vertices in interval 1, len2
//is the number of vertices in all previsous intervals (including interval 2).
//the format for indexForIndex[2*256:] is totUintNum(uint)(only include pairs),(start1,end1),(start2,end2),evid labels(16),recordPos(16).
// first, all same start are grouped together, in each group
//all end with same label are grouped. each number uses 5 bits, each time we at most process 16 pairs. we use one uint to represent use
//use record and isrecord.

uint* reduction(uint *edgelabelPartition,uint *edgeLabelPartition_dev, uint *vLabel_dev, uint*writeRowNum_dev, uint *writeRowNum,
               uint*baseAddr_dev,uint &stride,uint totSrcNum,uint intervalNum,Phase &curPhase,uint *partialEmb_dev,
               uint partialRowNum,bool isLastPhase, uint blocksize){

    uint *labelPart_dev = edgeLabelPartition_dev+edgelabelPartition[2];
    uint *neighborsData_dev = edgeLabelPartition_dev+edgelabelPartition[3];
    uint edgelabel = curPhase.edgeLabel;
    uint embLen = posOfVisitedV.size();
    uint64_t key = edgelabel;
    key = (key<<32) | 0xffffffff;

    /*uint *recordpos = recordPosData[key].addr;
    uint recordsize = recordPosData[key].size;
    uint *recordpos_dev = baseAddr_dev+stride;
    stride = stride+recordsize;
    cudaMemcpyAsync(recordpos_dev,recordpos,sizeof(uint)*recordsize,cudaMemcpyHostToDevice);*/
    std::vector<uint> &pairs = curPhase.pairs;
    sortReducePairs(pairs);
    if(pairs.size()>32){
        std::cout<<"not process this situation now"<<std::endl;
        return NULL;
    }
    uint totUintNum = (pairs.size()+5)/6;
    std::vector<uint> auxArray(1+6+16+16+1,0);
    auxArray[0]=totUintNum;
    uint groupsize = 0,tmp=0,index=1;
    for(uint i=0;i<pairs.size()/2;++i){
        uint svid = pairs[i*2];
        uint evid = pairs[i*2+1];
        //std::cout<<"after sort: svid="<<svid<<" evid="<<evid<<std::endl;
        uint evidlabel = queryLabel[evid];
        auxArray[1+6+i] = evidlabel;
        key = edgelabel;
        key = (key<<32)|0xffffffff;
        //uint *tmprecordpos = recordPosData[key].addr;
        //uint *tmprecordpos_dev = recordpos_dev+(tmprecordpos-recordpos);
        //auxArray[1+6+16+i] = tmprecordpos-recordpos;
        uint spos = posOfVisitedV[svid];
        uint epos = posOfVisitedV[evid];
        tmp = tmp | (epos<<groupsize);
        tmp = tmp | (spos<<(groupsize+5));
        groupsize = groupsize+10;
        if(groupsize==30){
            auxArray[index]=tmp;
            tmp=0;
            groupsize=0;
            ++index;
        }
    }
    if(tmp>0){
        auxArray[index] = tmp;
    }

    uint *aux_dev = baseAddr_dev+stride;
    stride = stride+1+6+16+16+1;
    cudaMemcpyAsync(aux_dev,auxArray.data(),sizeof(uint)*(1+6+16+16+1),cudaMemcpyHostToDevice);
    //cudaDeviceSynchronize();
    //cudaError_t err = cudaGetLastError();
    //std::cout<<"0#### "<<cudaGetErrorString(err)<<"#####partialRowNum="<<partialRowNum<<std::endl;
    std::cout<<"before reduction"<<std::endl;
    if(intervalNum<=1024){
        reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,partialRowNum,embLen,isLastPhase,0);
    }else{
        reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,partialRowNum,embLen,isLastPhase,blocksize);
    }
    //cudaDeviceSynchronize();
    //err = cudaGetLastError();
    //std::cout<<"1#### "<<cudaGetErrorString(err)<<std::endl;
    std::vector<uint*> tmppartResults;
    std::vector<uint> tmppartResultRowNum;
    cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
    partialRowNum = writeRowNum[0];
    std::cout<<"end reduction, partialRowNum="<<partialRowNum<<std::endl;
    if(partResults.size()>0){
        std::cout<<"there are more partial embedidngs to be processing in reduction"<<std::endl;
        uint allocnum = ((partialRowNum*embLen+31)>>5)<<5;
        uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
        cudaMemcpyAsync(tmpaddr,partialEmb_dev,sizeof(uint)*partialRowNum*embLen,cudaMemcpyDeviceToHost);
        tmppartResults.push_back(tmpaddr);
        tmppartResultRowNum.push_back(partialRowNum);
        writeRowNum[0] = 0;
        writeRowNum[1] = 0;
        cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
        uint totpartrownum=0;
        maxRowNum = (gpuAvailSpace-stride)/embLen;
        for(uint j=0;j<partResults.size();++j){
            tmpaddr = partResults[j];
            partialRowNum = partResultRowNum[j];
            totpartrownum+=partialRowNum;
            if(totpartrownum>=maxRowNum){
                totpartrownum = totpartrownum-partialRowNum;
                partialEmb_dev = baseAddr_dev+gpuAvailSpace-totpartrownum*embLen;
                if(intervalNum<=1024){
                    reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,totpartrownum,embLen,isLastPhase,0);
                }else{
                    reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,totpartrownum,embLen,isLastPhase,blocksize);
                }
                totpartrownum = partialRowNum;
                cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
                partialRowNum = writeRowNum[0];
                allocnum = ((partialRowNum*embLen+31)>>5)<<5;
                tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
                tmppartResults.push_back(tmpaddr);
                tmppartResultRowNum.push_back(partialRowNum);
                cudaMemcpyAsync(tmpaddr,partialEmb_dev,sizeof(uint)*partialRowNum*embLen,cudaMemcpyDeviceToHost);
                partialEmb_dev = baseAddr_dev+gpuAvailSpace-totpartrownum*embLen;
                tmpaddr = partResults[j];
                cudaMemcpyAsync(partialEmb_dev,tmpaddr,sizeof(uint)*totpartrownum*embLen,cudaMemcpyHostToDevice);
                writeRowNum[0] = 0;
                writeRowNum[1] = 0;
                cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            }else{
                partialEmb_dev = baseAddr_dev+gpuAvailSpace-totpartrownum*embLen;
                cudaMemcpyAsync(partialEmb_dev,tmpaddr,sizeof(uint)*partialRowNum*embLen,cudaMemcpyHostToDevice);
            }
        }
        partialEmb_dev = baseAddr_dev+gpuAvailSpace-totpartrownum*embLen;
        if(intervalNum<=1024){
            reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,totpartrownum,embLen,isLastPhase,0);
        }else{
            reductionPhase(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,neighborsData_dev,aux_dev,partialEmb_dev,intervalNum,totpartrownum,embLen,isLastPhase,blocksize);
        }
        cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
    }
    for(uint j=0;j<partResults.size();++j){
        uint *tmpaddr = partResults[j];
        free(tmpaddr);
    }
    partResultRowNum = tmppartResultRowNum;
    partResults = tmppartResults;
    //cudaMemcpyAsync(recordpos,recordpos_dev,sizeof(uint)*recordsize,cudaMemcpyDeviceToHost);
    //cudaDeviceSynchronize();
    //err = cudaGetLastError();
    //std::cout<<"2#### "<<cudaGetErrorString(err)<<std::endl;
    return partialEmb_dev;
}



int main(int argc, char *argv[]){
    cudaError_t err;
    cudaDeviceGetAttribute(&GPU_SM_NUM, cudaDevAttrMultiProcessorCount, 0);
    std::string inputFileName = std::string(argv[1]);//suffix .mygraph
    std::string metaFileName = std::string(argv[2]);//suffix .metadata
    std::string queryFileName = std::string(argv[3]);//suffix .query

	if(argc==5){
    	std::string checkFileName = std::string(argv[4]);//suffix .g;
    	loadDataBaseAndCheck(inputFileName,metaFileName,checkFileName);
    	return 0;
	}

    //std::cout<<"loading database"<<std::endl;
    allEdgeLabelPartitions = loadDataBase(inputFileName, metaFileName);

    //alloc gpu global memory, the unit is byte
    size_t total;
    uint *baseAddr_dev, stride=0, *partialEmb_dev;
    cudaMemGetInfo(&gpuAvailSpace,&total);
    //it seems that cudamemgetinfo triggers the bug
    gpuAvailSpace = ((gpuAvailSpace>>30)<<30);
    //std::cout<<"gpu avail mem:"<<gpuAvailSpace<<" gpu tot mem:"<<total<<std::endl;
    cudaMalloc((void**)&baseAddr_dev,gpuAvailSpace);
    //now gpuAvailSpace is how many uints available
    maxRowNum=0;
    gpuAvailSpace = gpuAvailSpace>>2;
    //err = cudaGetLastError();
    //std::cout<<"0 "<<cudaGetErrorString(err)<<std::endl;
    //alloc gpu memory for vLabel
    vLabel = allEdgeLabelPartitions[0xffffffff];
    uint totVNum = vLabel[0];
    uint *vLabel_dev = baseAddr_dev;
    stride = stride + totVNum+1;
    cudaMemcpyAsync(vLabel_dev,vLabel,sizeof(uint)*(totVNum+1),cudaMemcpyHostToDevice);

    //alloc gpu memory for totWriteRowNum
    //writeRowNum[0] is totWriteRowNum, writeRowNum[1] is extraWriteRowNum
    uint writeRowNum[2];
    uint *writeRowNum_dev = baseAddr_dev+stride;
    stride = stride + 2;

    //vLabel_dev and totWriteRowNum_dev are fixed
    baseAddr_dev = baseAddr_dev + stride;
    gpuAvailSpace = gpuAvailSpace-stride;
    stride = 0;
    //generate match order, and transfer the first edge label partition to device once we find the edge label of first phase
    //stride = stride + len;
    clock_t start, end;
    //std::cout<<"before generating matching order"<<std::endl;
    start = clock();
    genMatchOrder(queryFileName,matchOrder,repeatFlag,restricts,edgeLabelStat,queryLabel,allEdgeLabelPartitions,baseAddr_dev,stride);
    Phase &tmpPhase = matchOrder[0];
    std::cout<<"--------------restrictions--------------"<<std::endl;
    for(auto ite=restricts.begin();ite!=restricts.end();++ite){
        uint key = ite->first;
        uint value = ite->second;
        std::cout<<"restrict: "<<key<<"<"<<value<<std::endl;
    }
    std::cout<<"------------restrictions end------------"<<std::endl;
    std::cout<<"--------------phase order--------------"<<std::endl;
    for(uint i=0;i<matchOrder.size();++i){
        Phase &phase = matchOrder[i];
        std::cout << "phase "<<i<<" type "<<phase.phaseType<<" edge label:"<<phase.edgeLabel<<std::endl;
        std::vector<uint> &pairs = phase.pairs;
        for(uint j=0;j<pairs.size()/2;++j){
            std::cout<<"svid:"<<pairs[j*2]<<" evid:"<<pairs[j*2+1]<<" ";
        }
        std::cout<<std::endl;
    }
    std::cout<<"------------phase order end------------"<<std::endl;
    uint totWriteRowNum=0;
    uint label = tmpPhase.edgeLabel;
    uint *edgeLabelPartition = allEdgeLabelPartitions[label];
    uint len = edgeLabelPartition[0]-5;
    uint totSrcNum = edgeLabelPartition[1];
    uint intervalNum = edgeLabelPartition[2];
    uint blocksize = edgeLabelPartition[4];
    uint *edgeLabelPartition_dev = baseAddr_dev+stride-len;
    uint embLen;
    if(matchOrder[0].pairs.size()==2){
        embLen=2;
    }else{
        embLen=3;
    }
    writeRowNum[0] = 0;
    writeRowNum[1] = 0;
    cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
    //generate init embing
    std::cout<<"------------phase 0------------"<<std::endl;
    uint newstride = genInit(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,edgeLabelPartition_dev+edgeLabelPartition[3], 
        baseAddr_dev,stride,intervalNum,totSrcNum,false,blocksize);
    cudaDeviceSynchronize();
    err = cudaGetLastError();
    std::cout<<"2 "<<cudaGetErrorString(err)<<std::endl;   
    //get tot write row number
    cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
    uint partialRowNum = writeRowNum[0];

    //err = cudaGetLastError();
    //std::cout<<"2#### "<<cudaGetErrorString(err)<<std::endl;
    std::cout<<"after first init, partialRowNum="<<partialRowNum<<" maxRowNum="<<maxRowNum<<" totsrcnum="<<totSrcNum<<std::endl;
    std::ofstream outputfile;
    outputfile.open("partialEmbGPU_0.txt", std::ios::out);
    while(partialRowNum>=maxRowNum){
        //std::cout << "over size detect"<<std::endl;
        partialRowNum = partialRowNum-writeRowNum[1];
        std::cout<<"real partialrownum="<<partialRowNum<<std::endl;
        totWriteRowNum += partialRowNum;
        uint allocnum = ((partialRowNum*embLen+31)>>5)<<5;
        uint *tmpaddr = (uint*)malloc(sizeof(uint)*allocnum);
        cudaMemcpyAsync(tmpaddr,baseAddr_dev+newstride,sizeof(uint)*partialRowNum*embLen,cudaMemcpyDeviceToHost);
        cudaDeviceSynchronize();
        printPartialEmbGPU(tmpaddr,partialRowNum,embLen,outputfile);
        writeRowNum[0] = 0;
        writeRowNum[1] = 0;
        cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
        partResults.push_back(tmpaddr);
        partResultRowNum.push_back(partialRowNum);
        //std::cout<<"begin oversize init"<<std::endl;
        newstride = genInit(vLabel_dev,writeRowNum_dev,edgeLabelPartition_dev,edgeLabelPartition_dev+edgeLabelPartition[3], 
                            baseAddr_dev,stride,intervalNum,totSrcNum,true,blocksize);
        cudaMemcpy(writeRowNum,writeRowNum_dev,sizeof(uint)*2,cudaMemcpyDeviceToHost);
        partialRowNum = writeRowNum[0];
        std::cout<<"after oversize init, partialRowNum="<<partialRowNum<<" maxRowNum="<<maxRowNum<<std::endl;
    }
    stride = newstride;
    
    printPartialEmbGPU(baseAddr_dev+stride,partialRowNum,embLen,outputfile,1);
    outputfile.close();
    std::cout<<"------------phase 0 end------------"<<std::endl;
    writeRowNum[0] = writeRowNum[0]*embLen;
    partialEmb_dev = baseAddr_dev+gpuAvailSpace-writeRowNum[0];
    cudaDeviceSynchronize();
    err = cudaGetLastError();
    std::cout<<"2.5 "<<cudaGetErrorString(err)<<"####"<<writeRowNum[0]<<std::endl;
    copyPartialEmb(baseAddr_dev+stride,partialEmb_dev,writeRowNum[0]);
    cudaDeviceSynchronize();
    err = cudaGetLastError();
    std::cout<<"3 "<<cudaGetErrorString(err)<<std::endl;
    //printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,"partialEmbGPU_tmp.txt");
    //initialize totWriteRowNum_dev to 0
    writeRowNum[0] = 0;
    writeRowNum[1] = 0;
    cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
    //err = cudaGetLastError();
    //std::cout<<"3 "<<cudaGetErrorString(err)<<std::endl;
    uint preedgelabel = label,i;
    for(i=1;i<matchOrder.size();++i){
	    std::cout<<"-------------phase "<<i<<"-------------"<<std::endl;
        bool isLastPhase = i==(matchOrder.size()-1);
        stride = 0;
        Phase &curPhase = matchOrder[i];
        if(curPhase.edgeLabel!=preedgelabel){
            //std::cout<<"preedgelabel="<<preedgelabel<<" curedgelabel="<<curPhase.edgeLabel<<" totsize="<<allEdgeLabelPartitions.size()<<std::endl;
            preedgelabel = curPhase.edgeLabel;
            edgeLabelPartition = allEdgeLabelPartitions[preedgelabel];
            if(edgeLabelPartition==NULL){
                std::cout<<"partition is null"<<std::endl;
            }
            len = edgeLabelPartition[0]-5;
            totSrcNum = edgeLabelPartition[1];
            intervalNum = edgeLabelPartition[2];
            blocksize = edgeLabelPartition[4];
            //std::cout<<"len="<<len<<" totsrcNum="<<totSrcNum<<" intervalNum="<<intervalNum<<std::endl;
            cudaMemcpyAsync(edgeLabelPartition_dev,edgeLabelPartition+5,sizeof(uint)*len,cudaMemcpyHostToDevice);
        }
    	//cudaDeviceSynchronize();
        //err = cudaGetLastError();
        //std::cout<<"4 "<<i<<" "<<cudaGetErrorString(err)<<std::endl;
        stride = stride+len;
        uint recordrownum, recordemblen;
        if(curPhase.phaseType=='E'){
            if(curPhase.pairs.size()==2) { embLen += 1; }
            else { embLen += 2; }
            uint distance = (partialEmb_dev-baseAddr_dev)-stride;
            //std::cout<<"in Extension row num:"<<partialRowNum<<" it can fit in: "<<distance/4<<" rows"<<std::endl;
            /*if(i==6){
                std::ifstream tempfile("debugemb1.txt",std::ios::in|std::ios::binary);
                uint *partialEmb_host = (uint*)malloc(sizeof(uint)*recordrownum*recordemblen);
                tempfile.read((char*)partialEmb_host,sizeof(uint)*recordrownum*recordemblen);
                cudaMemcpy(partialEmb_dev,partialEmb_host,sizeof(uint)*recordrownum*recordemblen,cudaMemcpyHostToDevice);
                tempfile.close();
            }*/
            genExtEmb(vLabel_dev,writeRowNum_dev,writeRowNum,edgeLabelPartition_dev,edgeLabelPartition_dev+edgeLabelPartition[3],
                baseAddr_dev,stride,partialEmb_dev,intervalNum,totSrcNum,partialRowNum,curPhase,isLastPhase,false,blocksize);
            partialRowNum = writeRowNum[0];
		    if(partialRowNum==0){ break; }
            
            /*if(i==2){
                std::ofstream debug("debug1.txt",std::ios::out);
                printPartialEmbGPU(baseAddr_dev+stride,partialRowNum,embLen,debug,1);
                debug.close();
            }*/
            outputfile.open("partialEmbGPU_"+std::to_string(i)+".txt", std::ios::out);
            printPartialEmbGPU(baseAddr_dev+stride,partialRowNum,embLen,outputfile,1);
            for(int j=0;j<partResults.size();++j){
                uint *tmppart = partResults[j];
                uint tmppartrownum = partResultRowNum[j];
                printPartialEmbGPU(tmppart,tmppartrownum,embLen,outputfile);
            }
            outputfile.close();
            cudaDeviceSynchronize();
            err = cudaGetLastError();
            std::cout<<"7 "<<i<<" "<<cudaGetErrorString(err)<<" partialnum="<<partialRowNum<<std::endl;
            partialEmb_dev = baseAddr_dev+gpuAvailSpace-partialRowNum*embLen;
            //std::cout<<"test position:"<<partialRowNum<<" ="<<(partialEmb_dev-baseAddr_dev-stride)/embLen<<std::endl;
            copyPartialEmb(baseAddr_dev+stride,partialEmb_dev,partialRowNum*embLen);
            //std::string filename = "partialEmbGPU_ext_"+std::to_string(i)+".txt";
            //printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,filename);
            //printPartialEmbGPUDebug(partialEmb_dev,partialRowNum,embLen,filename);
            //err = cudaGetLastError();
            //std::cout<<"9 "<<i<<" "<<cudaGetErrorString(err)<<" partialnum="<<partialRowNum<<std::endl;
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            //err = cudaGetLastError();
            //std::cout<<"10 "<<i<<" "<<cudaGetErrorString(err)<<std::endl;
        }else{
            //std::cout<<"in Reduction"<<std::endl;
            partialEmb_dev = reduction(edgeLabelPartition,edgeLabelPartition_dev,vLabel_dev,writeRowNum_dev,writeRowNum,baseAddr_dev,stride,totSrcNum,intervalNum,curPhase,partialEmb_dev,partialRowNum,isLastPhase,blocksize);
            partialRowNum = writeRowNum[0];
		    if(partialRowNum==0){ break; }
            outputfile.open("partialEmbGPU_"+std::to_string(i)+".txt", std::ios::out);
            printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,outputfile,1);
            for(int j=0;j<partResults.size();++j){
                uint *tmppart = partResults[j];
                uint tmppartrownum = partResultRowNum[j];
                printPartialEmbGPU(tmppart,tmppartrownum,embLen,outputfile);
            }
            outputfile.close();
            //cudaDeviceSynchronize();
            //err = cudaGetLastError();
            //std::cout<<"10 "<<i<<" "<<cudaGetErrorString(err)<<" partialnum="<<partialRowNum<<std::endl;
		    //std::cout<<"reduction done:"<<partialRowNum<<std::endl;
            if(i<matchOrder.size()-1 && matchOrder[i+1].phaseType=='E'){
                uint *newpartialEmb_dev = baseAddr_dev+gpuAvailSpace-partialRowNum*embLen;
                copyPartialEmb(partialEmb_dev,newpartialEmb_dev,partialRowNum*embLen);
                partialEmb_dev = newpartialEmb_dev;
            }
            //std::cout<<"500 "<<i<<" "<<cudaGetErrorString(err)<<std::endl;
            //std::string filename = "partialEmbGPU_red_"+std::to_string(i)+".txt";
            //std::cout<<"rownum="<<partialRowNum<<std::endl;
            //printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,filename);
            //printPartialEmbGPUDebug(partialEmb_dev,partialRowNum,embLen,filename);
            writeRowNum[0] = 0;
            writeRowNum[1] = 0;
            cudaMemcpyAsync(writeRowNum_dev,writeRowNum,sizeof(uint)*2,cudaMemcpyHostToDevice);
            //err = cudaGetLastError();
            //std::cout<<"11 "<<i<<" "<<cudaGetErrorString(err)<<" partialnum="<<partialRowNum<<std::endl;
        }
	    std::cout<<"-------------phase "<<i<<" end-------------"<<std::endl;
    }
    cudaDeviceSynchronize();
    end = clock();
    double endtime=(double)(end-start)/CLOCKS_PER_SEC;
    std::cout<<"time: "<<endtime*1000<<std::endl;
    err = cudaGetLastError();
    std::cout<<cudaGetErrorString(err)<<std::endl;
    std::cout<<"Result: "<<partialRowNum<<" extra size"<<partResults.size()<<std::endl;
    std::cout<<"PartialRowNum: "<<partialRowNum<<" embLen:  "<<embLen<<std::endl;
    outputfile.open("partialResultsGPU.txt", std::ios::out);
    printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,outputfile,1);
    for(i=0;i<partResults.size();i++){
        partialRowNum = partResultRowNum[i];
        uint *tmppart = partResults[i];
        printPartialEmbGPU(tmppart,partialRowNum,embLen,outputfile);
    }
    outputfile.close();
    //printPartialEmbGPU(partialEmb_dev,partialRowNum,embLen,"partialResultsGPU.txt");
    return 0;
}
